#include <cutil_inline.h>
#include <device_launch_parameters.h>
#include <thrust\sort.h>
#include <thrust\device_vector.h>
#include <thrust\device_ptr.h>
#include <thrust\functional.h>
#include <thrust\transform.h>
#include "config.cuh"

typedef int test_t;

__global__ void init_kernel(int *index, int *init, int *baseId, int *no, int *etype, 
	float *time, float *duration, float *speed, float *pos, int N)
{
	unsigned int i = threadIdx.x+blockIdx.x*blockDim.x;
	if(i<N){
		int eidx = init[i];
		int vno = no[eidx];
		float vtime = time[eidx];
		float vduration = duration[eidx];
		float vspeed = speed[eidx];
		float vpos = pos[eidx];
		index[eidx]=eidx;

		float termiPos = vpos+vduration*vspeed/3600;
		if(termiPos>40)
			termiPos = 40;
		float currPos = vpos;
		float nextPos = ((int)(currPos/2)+1)*2;
		while(nextPos<termiPos){
			eidx++;
			index[eidx] = eidx;
			time[eidx] = vtime = vtime+(nextPos-currPos)/vspeed*3600;
			speed[eidx] = vspeed;
			pos[eidx] = nextPos;
			no[eidx] = vno;
			duration[eidx] = 0;
			etype[eidx] = TERMI;
			baseId[eidx] = currPos/2;

			eidx++;
			index[eidx] = eidx;
			time[eidx] = vtime;
			speed[eidx] = vspeed;
			pos[eidx] = nextPos;
			no[eidx] = vno;
			duration[eidx] = vduration-(nextPos-vpos)/vspeed*3600;
			etype[eidx] = HANDO;
			baseId[eidx] = nextPos/2;

			currPos = nextPos;
			nextPos = nextPos+2;
		}
		eidx++;
		index[eidx] = eidx;
		time[eidx] = vtime+(termiPos-currPos)/vspeed*3600;
		speed[eidx] = vspeed;
		pos[eidx] = termiPos;
		no[eidx] = vno;
		duration[eidx] = 0;
		etype[eidx] = TERMI;
		baseId[eidx] = currPos/2;
	}
}

//float *d_init, float *d_duration, float *d_speed, float *d_pos can be freed
extern "C" void launch_init_kernel(da_t da, int n)
{
	//n = SAMPLENO;
	init_kernel<<<1+n/BLOCKSIZE,BLOCKSIZE>>>(da.index, da.init, da.baseId, da.no, da.etype, 
		da.time, da.duration, da.speed, da.pos, n);
	cudaFree(da.speed);
	cudaFree(da.pos);
	cudaFree(da.duration);
	//cudaFree(da.init);
}

typedef thrust::tuple<int, float, int> KeyType; //base, time, no
struct Comp : public thrust::binary_function<KeyType, KeyType, bool>
{
	__host__ __device__
		bool operator () (const KeyType& a, const KeyType& b) const
	{
		if(thrust::get<0>(a)!=thrust::get<0>(b))
			return thrust::get<0>(a) < thrust::get<0>(b);
		else if(thrust::get<1>(a)!=thrust::get<1>(b))
			return thrust::get<1>(a) < thrust::get<1>(b);
		else
			return thrust::get<2>(a) < thrust::get<2>(b);
	}
};

//float *d_time can be freed
//int *index, int *baseId and float *time are sorted
//int *no, int *etype and int *sum are not sorted
extern "C" void launch_sort_kernel(da_t da, int n)
{
	int *tno;
	cudaMalloc((void**)&tno, n*sizeof(int));
	cudaMemcpy(tno, da.no, n*sizeof(int), cudaMemcpyDeviceToDevice);

	thrust::device_ptr<int> index_ptr(da.index);
	thrust::device_ptr<int> baseId_ptr(da.baseId);
	thrust::device_ptr<int> no_ptr(tno);
	thrust::device_ptr<float> time_ptr(da.time);

	typedef thrust::device_vector<int>::iterator IntIter;
	typedef thrust::device_vector<float>::iterator FloatIter;
	typedef thrust::tuple<IntIter, FloatIter, IntIter> KeyIterTuple; //base, time, no
	typedef thrust::tuple<IntIter> ValIterTuple; //index
	typedef thrust::zip_iterator<KeyIterTuple> KeyIter;
	typedef thrust::zip_iterator<ValIterTuple> ValIter;

	KeyIter key_begin(thrust::make_tuple(baseId_ptr, time_ptr, no_ptr));
	KeyIter key_end(thrust::make_tuple(baseId_ptr+n, time_ptr+n, no_ptr+n));
	ValIter val_begin(thrust::make_tuple(index_ptr));

	thrust::sort_by_key(key_begin, key_end, val_begin, Comp());	
	//cudaFree(da.time);
}

inline __device__ int scan_warp(volatile int* data, volatile int* flag, const int tid){
	const unsigned int lane = tid&31;
	if(lane >= 1) if(flag[tid] == flag[tid-1]) data[tid] = data[tid-1] + data[tid];
	if(lane >= 2) if(flag[tid] == flag[tid-2]) data[tid] = data[tid-2] + data[tid];
	if(lane >= 4) if(flag[tid] == flag[tid-4]) data[tid] = data[tid-4] + data[tid];
	if(lane >= 8) if(flag[tid] == flag[tid-8]) data[tid] = data[tid-8] + data[tid];
	if(lane >= 16) if(flag[tid] == flag[tid-16]) data[tid] = data[tid-16] + data[tid];
	return data[tid];
}

inline __device__ int scan_block(volatile int* data, volatile int* flag, const int tid){
	const unsigned int lane = tid & 31;
	const unsigned int wid	= tid >> 5; //warp id

	int vdata = scan_warp(data, flag, tid);
	int vflag = flag[tid];
	__syncthreads();
	if(lane == 31 && tid & (BLOCKSIZE-1) != BLOCKSIZE-1) {
		data[wid] = data[tid]; flag[wid] = flag[tid];}
	__syncthreads();
	if(tid & (BLOCKSIZE-1) == BLOCKSIZE-1) {
		data[wid] = data[tid]; flag[wid] = flag[tid];}
	if(wid == 0)
		scan_warp(data, flag, tid);
	__syncthreads();
	if(wid>0) 
		if(vflag == flag[wid-1]) vdata = data[wid-1]+vdata;
	__syncthreads();
	data[tid] = vdata; 
	__syncthreads();

	return vdata;
}

extern __shared__ int smem[];
__global__ void scan_kernel(int *index, int *data, int *flag, int *tdata, int *tflag, int elemno){
	const unsigned int tid	= threadIdx.x;
	const unsigned int i	= threadIdx.x+blockDim.x*blockIdx.x;
	int *sdata = (int*)smem;
	int *sflag = (int*)&(smem[BLOCKSIZE]);

	if(i < elemno){
		if(index != NULL){
			sdata[tid] = data[index[i]];
			sflag[tid] = flag[i];
		} else {
			sdata[tid] = data[i];
			sflag[tid] = flag[i];
		}
	}
	__syncthreads();
	scan_block(sdata, sflag, tid);
	//scan_warp(sdata, sflag, tid);
	if(tid == 0 && tdata != NULL){
		tdata[blockIdx.x+1] = sdata[BLOCKSIZE-1];
		tflag[blockIdx.x+1] = sflag[BLOCKSIZE-1];
	}
	if(i == 0 && tdata != NULL){
		tdata[0] = 0;
		tflag[0] = 0;
	}
	if(i < elemno){
		if(index!=NULL)
			data[index[i]] = sdata[tid];
		else
			data[i]=sdata[tid];
	}
}

__global__ void addback_kernel(int* detect, int *index, int *data, int *flag, int *tdata, int *tflag, int elemno){
	const unsigned int tid	= threadIdx.x;
	const unsigned int i	= threadIdx.x+blockDim.x*blockIdx.x;

	int idx = -1;
	if(i<elemno)
		idx = index[i];
	if(tid==0) {
		smem[0] = tdata[blockIdx.x]; //sdata
		smem[1] = tflag[blockIdx.x]; //sflag
	}
	__syncthreads();
	if(i<elemno)
		if(flag[i] == smem[1])
			data[idx] = data[idx]+smem[0];
	if(i==0)
		detect[i]=FALSE;
}

int iter=0;
char iterinchar[25];
extern "C" void launch_scan_kernel(da_t da, int n){
	int *tdata;
	int *tflag;
	int gsize = GSIZE(n);
	cudaMalloc((void**)&tdata, gsize*sizeof(int));
	cudaMalloc((void**)&tflag, gsize*sizeof(int));

	cudaMemcpy(da.sum, da.etype, n*sizeof(int), cudaMemcpyDeviceToDevice);
	scan_kernel<<<gsize, BLOCKSIZE, SMEMSIZE>>>(da.index, da.sum, da.baseId, tdata, tflag, n);
	int tempgsize = GSIZE(gsize);
	scan_kernel<<<tempgsize, BLOCKSIZE, SMEMSIZE>>>(NULL, tdata, tflag, NULL, NULL, gsize);
	int smemsize = sizeof(int)+sizeof(int); //allocate two int for flag and data in tdata and tflag.
	addback_kernel<<<gsize, BLOCKSIZE, smemsize>>>(da.detect, da.index, da.sum, da.baseId, tdata, tflag, n);

	//sprintf(iterinchar, "scan%d.txt", iter);
	//printinfo(iterinchar);
}

__global__ void corr_kernel1(int *etype, int *sum, int *no, int* detect, int n){
	const unsigned int tid	= threadIdx.x;
	const unsigned int i	= threadIdx.x + blockDim .x * blockIdx.x;
	int tno = 0;
	if(sum[i] == 11 && etype[i] == 1){
		//if(sum[i] > 10 && etype[i] == 1){
		int idx = i+1;
		etype[i] = 0;
		tno = no[i];
		no[i] = -no[i];
		//if(no[idx] == tno || no[idx] == -tno){
		while(no[idx] == tno || no[idx] == -tno){
			etype[idx] = 0;
			idx++;
		}
		detect[0] = TRUE;
	}
}

__global__ void corr_kernel2(int *etype, int *sum, int *no, int* detect, int n){
	const unsigned int tid	= threadIdx.x;
	const unsigned int i	= threadIdx.x + blockDim .x * blockIdx.x;
	int tno = 0;
	if(no[i] < 0 && sum[i] < 10){
		char t = 0;
		int idx = i+1;
		etype[i]=1;
		tno = no[i];
		no[i] = -no[i];
		while(no[idx] == tno || no[idx] == -tno){
			etype[idx] = t%2 == 0 ? -1 : 1;
			idx++;
			t++;
		}
		detect[0] = TRUE;
	}
}

__global__ void corr_kernel2_1(int *init, int *etype, int *sum, int *no, int *detect, int n){
	const unsigned int tid	= threadIdx.x;
	const unsigned int i	= threadIdx.x + blockDim .x * blockIdx.x;
	if(i<n){
		int idx = init[i];
		int flag = FALSE;
		while(no[idx] == no[idx+1] || no[idx] == -no[idx+1] ){
			if(flag == TRUE){
				etype[idx] = 0;
				if(no[idx]<0)
					no[idx]=-no[idx];
			}
			if(flag == FALSE && sum[idx] == 11 && etype[idx] == 1){
				flag = TRUE;
				no[idx] = -no[idx];
				etype[idx] = 0;
				detect[0]=TRUE;
			}
			idx++;
		}
		if(flag == TRUE)
			etype[idx]=0;
	}
}

__global__ void corr_kernel2_2(int *init, int *etype, int *sum, int *no, int *detect, int n){
	const unsigned int tid	= threadIdx.x;
	const unsigned int i	= threadIdx.x + blockDim .x * blockIdx.x;
	if(i<n){
		int idx = init[i];
		int flag = FALSE;
		int t = 0;
		while(no[idx] == no[idx+1] || no[idx] == -no[idx+1]){
			if(sum[idx] < 10 && no[idx] < 0){
				flag = TRUE;
				no[idx] = -no[idx];
				detect[0] = TRUE;
			}
			if(flag == TRUE){
				etype[idx]=(t%2==0?1:-1);
				t++;
			}
			idx++;
		}
		if(flag == TRUE)
			etype[idx]=-1;
	}
}

extern "C" bool launch_corr_kernel(da_t da, int n){
	int gsize = GSIZE(n);
#if VERSION == 0	
	corr_kernel1<<<gsize, BLOCKSIZE>>>(da.etype, da.sum, da.no, da.detect, n);
	corr_kernel2<<<gsize, BLOCKSIZE>>>(da.etype, da.sum, da.no, da.detect, n);
#else
	corr_kernel2_1<<<gsize, BLOCKSIZE>>>(da.init, da.etype, da.sum, da.no, da.detect, n);
	corr_kernel2_2<<<gsize, BLOCKSIZE>>>(da.init, da.etype, da.sum, da.no, da.detect, n);
#endif
	int detect = 0;
	cudaMemcpy(&detect, da.detect, sizeof(int), cudaMemcpyDeviceToHost);
	return detect == TRUE;
}
